#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <iostream>
#include <CL/cl.h>
#include <CL/opencl.h>
#include <sys/time.h>
#include <unistd.h>

using namespace std;

#define errorMessage(X) X

int main() {
  //first section
  cl_int error = 0;   // Used to handle error codes
  cl_uint numPlatforms;
  cl_platform_id platforms[10];
  cl_platform_id platform;
  cl_context context;
  cl_command_queue queue;
  cl_device_id device;

  // All platform
  error = clGetPlatformIDs(10, platforms, &numPlatforms);


  // Platform
  if (error != CL_SUCCESS) {
     cout << "Error getting platform id: " << errorMessage(error) << endl;
     exit(error);
  }

  if ( numPlatforms < 2 ) {
      cout << "No OpenCL platforms" << endl;
      exit(1);
  }

  //choose the first one
  platform = platforms[0];

  // Device
  error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
  if (error != CL_SUCCESS) {
     cout << "Error getting device ids: " << errorMessage(error) << endl;
     exit(error);
  }
  // Context
  context = clCreateContext(0, 1, &device, NULL, NULL, &error);
  if (error != CL_SUCCESS) {
     cout << "Error creating context: " << errorMessage(error) << endl;
     exit(error);
  }
  // Command-queue
  queue = clCreateCommandQueue(context, device, 0, &error);
  if (error != CL_SUCCESS) {
     cout << "Error creating command queue: " << errorMessage(error) << endl;
     exit(error);
  }


  const int size = 1024*1024*32;
  float* src_a_h = new float[size];
  float* src_b_h = new float[size];
  float* res_h = new float[size];
  // Initialize both vectors
  for (int i = 0; i < size; i++) {
     src_a_h[i] = src_b_h[i] = (float)i;
  }

  const int mem_size = sizeof(float)*size;
  // Allocates a buffer of size mem_size and copies mem_size bytes from src_a_h
  cl_mem src_a_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_a_h, &error);
  cl_mem src_b_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_b_h, &error);
//  cl_mem src_a_d = clCreateBuffer(context, CL_MEM_READ_ONLY , mem_size, NULL, &error);
//  cl_mem src_b_d = clCreateBuffer(context, CL_MEM_READ_ONLY, mem_size, NULL, &error);

  cl_mem res_d = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size, NULL, &error);


  // Creates the program
  // Uses NVIDIA helper functions to get the code string and it's size (in bytes)
  size_t src_size = 0;
//  const char* path = shrFindFilePath("vector_add_gpu.cl", NULL);
//  const char* source = oclLoadProgSource(path, "", &src_size);
  const char *ocl_test_program[] = {\
      "__kernel void vector_add_gpu (__global const float* src_a, __global const float* src_b, __global float* res, const int num)"\
      "{"\
      "  const int idx = get_global_id(0);"\
      "  if (idx < num) res[idx] = src_a[idx] + src_b[idx];"\
      "}"
  };


  cl_program program = clCreateProgramWithSource(context, 1, ocl_test_program, &src_size, &error);
  assert(error == CL_SUCCESS);

  // Builds the program
  error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
  assert(error == CL_SUCCESS);

  // Shows the log
  char* build_log;
  size_t log_size;
  // First call to know the proper size
  clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
  build_log = new char[log_size+1];
  // Second call to get the log
  clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);
  build_log[log_size] = '\0';
  cout << "Build Log" << build_log << endl;
  delete[] build_log;

  // Extracting the kernel
  cl_kernel vector_add_kernel = clCreateKernel(program, "vector_add_gpu", &error);
  assert(error == CL_SUCCESS);


  // Enqueuing parameters
  // Note that we inform the size of the cl_mem object, not the size of the memory pointed by it
  error = clSetKernelArg(vector_add_kernel, 0, sizeof(cl_mem), &src_a_d);
  error |= clSetKernelArg(vector_add_kernel, 1, sizeof(cl_mem), &src_b_d);
  error |= clSetKernelArg(vector_add_kernel, 2, sizeof(cl_mem), &res_d);
  error |= clSetKernelArg(vector_add_kernel, 3, sizeof(int), &size);
  assert(error == CL_SUCCESS);



  //slight change, and cp here
//  error = clEnqueueWriteBuffer(queue, src_a_d, CL_FALSE, 0, mem_size, src_a_h, 0, NULL, NULL);
//  error |= clEnqueueWriteBuffer(queue, src_a_d, CL_FALSE, 0, mem_size, src_b_h, 0, NULL, NULL);


  // Launching kernel
  const size_t local_ws = 512;    // Number of work-items per work-group
  // shrRoundUp returns the smallest multiple of local_ws bigger than size
  const size_t global_ws = 512; //shrRoundUp(local_ws, size);    // Total number of work-items
  struct timeval start, end;
  long mtime, seconds, useconds;
  gettimeofday(&start, NULL);
  error = clEnqueueNDRangeKernel(queue, vector_add_kernel, 1, NULL, &global_ws, &local_ws, 0, NULL, NULL);
  gettimeofday(&end, NULL);
  seconds  = end.tv_sec  - start.tv_sec;
  useconds = end.tv_usec - start.tv_usec;
  mtime = ((seconds) * 1000 + useconds/1000.0) + 0.5;

  printf("===>Time used %d ms\n", mtime);
  assert(error == CL_SUCCESS);


  float* check = new float[size];
  clEnqueueReadBuffer(queue, res_d, CL_TRUE, 0, mem_size, check, 0, NULL, NULL);

//  for(unsigned i = 0; i < size; i ++){
//      cout << check[i] << " ";
//      if(i % 16 == 15){ cout << endl;}
//  }

  // Cleaning up
  delete[] src_a_h;
  delete[] src_b_h;
  delete[] res_h;
  delete[] check;
  clReleaseKernel(vector_add_kernel);
  clReleaseCommandQueue(queue);
  clReleaseContext(context);
  clReleaseMemObject(src_a_d);
  clReleaseMemObject(src_b_d);
  clReleaseMemObject(res_d);

}
